[NVPTX] Add prefetch tensormap variant #146203

abhilash1910 · 2025-06-28T08:30:18Z

[NVPTX] Add Prefetch tensormap intrinsics
This PR adds prefetch intrinsics with the relevant tensormap_space.

Lit tests are added as part of prefetch.ll
The generated PTX is verified with a 12.3 ptxas executable.
Added docs for these intrinsics in NVPTXUsage.rst.

For more information, refer to the PTX ISA for prefetch intrinsic : Prefetch Tensormap

@durga4github @schwarzschild-radius

llvmbot · 2025-06-28T08:30:46Z

@llvm/pr-subscribers-llvm-ir

Author: Abhilash Majumder (abhilash1910)

Changes

[NVPTX] Add Prefetch tensormap intrinsics
This PR adds prefetch intrinsics with the relevant tensormap_space.

Lit tests are added as part of prefetch.ll
The generated PTX is verified with a 12.3 ptxas executable.
Added docs for these intrinsics in NVPTXUsage.rst.

For more information, refer to the PTX ISA for prefetch intrinsic : Prefetch Tensormap

@durga4github @schwarzschild-radius

Full diff: https://github.com/llvm/llvm-project/pull/146203.diff

4 Files Affected:

(modified) llvm/docs/NVPTXUsage.rst (+7-1)
(modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+4)
(modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+9)
(modified) llvm/test/CodeGen/NVPTX/prefetch.ll (+30)

diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 11017fe4e01b4..ca951811b73dd 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -971,6 +971,9 @@ Syntax:
   declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)
   declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)
   
+  declare void  @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+  declare void  @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+  
   declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
   declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
 
@@ -983,7 +986,10 @@ The '``@llvm.nvvm.prefetch.*``' and '``@llvm.nvvm.prefetchu.*``' intrinsic
 correspond to the '``prefetch.*``;' and '``prefetchu.*``' family of PTX instructions. 
 The '``prefetch.*``' instructions bring the cache line containing the
 specified address in global or local memory address space into the 
-specified cache level (L1 or L2). The '`prefetchu.*``' instruction brings the cache line 
+specified cache level (L1 or L2). If the '``.tensormap``' qualifier is specified then the 
+prefetch instruction brings the cache line containing the specified address in the 
+'``.const``' or '``.param memory``' state space for subsequent use by the '``cp.async.bulk.tensor``' 
+instruction.The '`prefetchu.*``' instruction brings the cache line 
 containing the specified generic address into the specified uniform cache level.
 If no address space is specified, it is assumed to be generic address. The intrinsic 
 uses and eviction priority which can be accessed by the '``.level::eviction_priority``' modifier.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 0375f29ad8906..0678bba51e4a3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -137,6 +137,7 @@
 
 def llvm_global_ptr_ty  : LLVMQualPointerType<1>;         // (global)ptr
 def llvm_shared_ptr_ty  : LLVMQualPointerType<3>;         // (shared)ptr
+def llvm_constant_ptr_ty: LLVMQualPointerType<4>;         // (const)ptr
 def llvm_local_ptr_ty   : LLVMQualPointerType<5>;         // (local)ptr
 def llvm_tmem_ptr_ty    : LLVMQualPointerType<6>;         // (tensor memory)ptr
 def llvm_shared_cluster_ptr_ty : LLVMQualPointerType<7>;  // (shared_cluster)ptr
@@ -2092,6 +2093,9 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
     def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
   }
 
+  def int_nvvm_prefetch_tensormap_p0: DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+  def int_nvvm_prefetch_tensormap_p4 : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
+  
   foreach eviction_priority = ["evict_normal", "evict_last"] in
     def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
 
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cc1fd027d8515..8afc7063c363a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -760,6 +760,15 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
 def PREFETCH_LOCAL_L1  : PREFETCH_INTRS<"prefetch.local.L1">;
 def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
 def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
+def PREFETCH_CONST_TENSORMAP        : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+                                      "prefetch.const.tensormap",
+                                      [(int_nvvm_prefetch_tensormap_p4 addr:$addr)]>,
+                                      Requires<[hasPTX<80>, hasSM<90>]>;
+                                      
+def PREFETCH_GENERIC_TENSORMAP      : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+                                      "prefetch.tensormap",
+                                      [(int_nvvm_prefetch_tensormap_p0 addr:$addr)]>,
+                                      Requires<[hasPTX<80>, hasSM<90>]>;
 
 def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
                                       "prefetch.global.L2::evict_normal",
diff --git a/llvm/test/CodeGen/NVPTX/prefetch.ll b/llvm/test/CodeGen/NVPTX/prefetch.ll
index a64e4fe7a508e..b63155ff49185 100644
--- a/llvm/test/CodeGen/NVPTX/prefetch.ll
+++ b/llvm/test/CodeGen/NVPTX/prefetch.ll
@@ -12,6 +12,9 @@ declare void  @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
 declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)
 declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)
 
+declare void  @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+declare void  @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+
 declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
 declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
 
@@ -67,6 +70,33 @@ define void @prefetch_(ptr %ptr) {
   ret void
 }
 
+
+define void @prefetch_generic_tensormap(ptr %ptr) {
+; CHECK-PTX64-LABEL: prefetch_generic_tensormap(
+; CHECK-PTX64:       {
+; CHECK-PTX64-NEXT:    .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT:  // %bb.0:
+; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0];
+; CHECK-PTX64-NEXT:    prefetch.tensormap [%rd1];
+; CHECK-PTX64-NEXT:    ret;
+  tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+  ret void
+}
+
+define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
+; CHECK-PTX64-LABEL: prefetch_const_tensormap(
+; CHECK-PTX64:       {
+; CHECK-PTX64-NEXT:    .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT:  // %bb.0:
+; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
+; CHECK-PTX64-NEXT:    prefetch.const.tensormap [%rd1];
+; CHECK-PTX64-NEXT:    ret;
+  tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4)  %const_ptr)
+  ret void
+}
+
 define void @prefetchu_l1(ptr %ptr) {
 ; CHECK-PTX64-LABEL: prefetchu_l1(
 ; CHECK-PTX64:       {

llvmbot · 2025-06-28T08:30:47Z

@llvm/pr-subscribers-backend-nvptx

Author: Abhilash Majumder (abhilash1910)

Changes

[NVPTX] Add Prefetch tensormap intrinsics
This PR adds prefetch intrinsics with the relevant tensormap_space.

Lit tests are added as part of prefetch.ll
The generated PTX is verified with a 12.3 ptxas executable.
Added docs for these intrinsics in NVPTXUsage.rst.

For more information, refer to the PTX ISA for prefetch intrinsic : Prefetch Tensormap

@durga4github @schwarzschild-radius

Full diff: https://github.com/llvm/llvm-project/pull/146203.diff

4 Files Affected:

(modified) llvm/docs/NVPTXUsage.rst (+7-1)
(modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+4)
(modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+9)
(modified) llvm/test/CodeGen/NVPTX/prefetch.ll (+30)

diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 11017fe4e01b4..ca951811b73dd 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -971,6 +971,9 @@ Syntax:
   declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)
   declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)
   
+  declare void  @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+  declare void  @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+  
   declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
   declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
 
@@ -983,7 +986,10 @@ The '``@llvm.nvvm.prefetch.*``' and '``@llvm.nvvm.prefetchu.*``' intrinsic
 correspond to the '``prefetch.*``;' and '``prefetchu.*``' family of PTX instructions. 
 The '``prefetch.*``' instructions bring the cache line containing the
 specified address in global or local memory address space into the 
-specified cache level (L1 or L2). The '`prefetchu.*``' instruction brings the cache line 
+specified cache level (L1 or L2). If the '``.tensormap``' qualifier is specified then the 
+prefetch instruction brings the cache line containing the specified address in the 
+'``.const``' or '``.param memory``' state space for subsequent use by the '``cp.async.bulk.tensor``' 
+instruction.The '`prefetchu.*``' instruction brings the cache line 
 containing the specified generic address into the specified uniform cache level.
 If no address space is specified, it is assumed to be generic address. The intrinsic 
 uses and eviction priority which can be accessed by the '``.level::eviction_priority``' modifier.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 0375f29ad8906..0678bba51e4a3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -137,6 +137,7 @@
 
 def llvm_global_ptr_ty  : LLVMQualPointerType<1>;         // (global)ptr
 def llvm_shared_ptr_ty  : LLVMQualPointerType<3>;         // (shared)ptr
+def llvm_constant_ptr_ty: LLVMQualPointerType<4>;         // (const)ptr
 def llvm_local_ptr_ty   : LLVMQualPointerType<5>;         // (local)ptr
 def llvm_tmem_ptr_ty    : LLVMQualPointerType<6>;         // (tensor memory)ptr
 def llvm_shared_cluster_ptr_ty : LLVMQualPointerType<7>;  // (shared_cluster)ptr
@@ -2092,6 +2093,9 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
     def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
   }
 
+  def int_nvvm_prefetch_tensormap_p0: DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+  def int_nvvm_prefetch_tensormap_p4 : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
+  
   foreach eviction_priority = ["evict_normal", "evict_last"] in
     def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
 
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cc1fd027d8515..8afc7063c363a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -760,6 +760,15 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
 def PREFETCH_LOCAL_L1  : PREFETCH_INTRS<"prefetch.local.L1">;
 def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
 def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
+def PREFETCH_CONST_TENSORMAP        : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+                                      "prefetch.const.tensormap",
+                                      [(int_nvvm_prefetch_tensormap_p4 addr:$addr)]>,
+                                      Requires<[hasPTX<80>, hasSM<90>]>;
+                                      
+def PREFETCH_GENERIC_TENSORMAP      : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+                                      "prefetch.tensormap",
+                                      [(int_nvvm_prefetch_tensormap_p0 addr:$addr)]>,
+                                      Requires<[hasPTX<80>, hasSM<90>]>;
 
 def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
                                       "prefetch.global.L2::evict_normal",
diff --git a/llvm/test/CodeGen/NVPTX/prefetch.ll b/llvm/test/CodeGen/NVPTX/prefetch.ll
index a64e4fe7a508e..b63155ff49185 100644
--- a/llvm/test/CodeGen/NVPTX/prefetch.ll
+++ b/llvm/test/CodeGen/NVPTX/prefetch.ll
@@ -12,6 +12,9 @@ declare void  @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
 declare void  @llvm.nvvm.prefetch.L1(ptr %ptr)
 declare void  @llvm.nvvm.prefetch.L2(ptr %ptr)
 
+declare void  @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+declare void  @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+
 declare void  @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
 declare void  @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
 
@@ -67,6 +70,33 @@ define void @prefetch_(ptr %ptr) {
   ret void
 }
 
+
+define void @prefetch_generic_tensormap(ptr %ptr) {
+; CHECK-PTX64-LABEL: prefetch_generic_tensormap(
+; CHECK-PTX64:       {
+; CHECK-PTX64-NEXT:    .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT:  // %bb.0:
+; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0];
+; CHECK-PTX64-NEXT:    prefetch.tensormap [%rd1];
+; CHECK-PTX64-NEXT:    ret;
+  tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+  ret void
+}
+
+define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
+; CHECK-PTX64-LABEL: prefetch_const_tensormap(
+; CHECK-PTX64:       {
+; CHECK-PTX64-NEXT:    .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT:  // %bb.0:
+; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
+; CHECK-PTX64-NEXT:    prefetch.const.tensormap [%rd1];
+; CHECK-PTX64-NEXT:    ret;
+  tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4)  %const_ptr)
+  ret void
+}
+
 define void @prefetchu_l1(ptr %ptr) {
 ; CHECK-PTX64-LABEL: prefetchu_l1(
 ; CHECK-PTX64:       {

llvm/docs/NVPTXUsage.rst

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

durga4github

The latest revision lgtm

llvm/include/llvm/IR/IntrinsicsNVVM.td

github-actions · 2025-07-07T09:12:12Z

✅ With the latest revision this PR passed the C/C++ code formatter.

AlexMaclean · 2025-07-07T15:24:36Z

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

          BasicNVPTXInst<(outs), (ins ADDR:$addr),
          InstName,
-          [(!cast<Intrinsic>(!strconcat("int_nvvm_",
-          !subst(".", "_", InstName))) addr:$addr)]>,
-          Requires<[hasPTX<80>, hasSM<90>]>;


Why have these predicates been removed?

Yes this needs to be added, fixing in progress. Thanks

Has this been fixed?

Missed this, will update. thanks

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

AlexMaclean · 2025-07-07T15:27:29Z

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

@@ -587,6 +588,9 @@ Value *NVPTXTTIImpl::rewriteIntrinsicWithAddressSpace(IntrinsicInst *II,
      return ConstantInt::get(II->getType(), *R);
    return nullptr;
  }
+  case Intrinsic::nvvm_prefetch_tensormap: {
+    return nullptr;


In order for this to be useful we need to actually rewrite the intrinsic here. That is, create an nvvm_prefetch_tensormap intrinsic with NewV. It would also be good to add a couple tests to confirm this works as expected.

AlexMaclean · 2025-07-09T19:59:31Z

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

+    Module *M = II->getModule();
+    Function *NewIntrinsic = Intrinsic::getDeclaration(
+        M, Intrinsic::nvvm_prefetch_tensormap, {NewV->getType()});
+    CallInst *NewCall = Builder.CreateCall(NewIntrinsic, {NewV});
+    NewCall->setAttributes(II->getAttributes());


Use Builder.CreateUnaryIntrinsic

AlexMaclean · 2025-07-09T20:01:30Z

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

          BasicNVPTXInst<(outs), (ins ADDR:$addr),
          InstName,
-          [(!cast<Intrinsic>(!strconcat("int_nvvm_",
-          !subst(".", "_", InstName))) addr:$addr)]>,
-          Requires<[hasPTX<80>, hasSM<90>]>;


Has this been fixed?

AlexMaclean · 2025-07-09T20:03:51Z

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

+  def "" : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+           !strconcat("prefetch", !if(!eq(addrspace_name, ""), "",
+           !strconcat(".", addrspace_name)), ".tensormap [$addr];"),
+           [(pattern_frag addr:$addr)]>,
+           Requires<[hasPTX<80>, hasSM<90>]>;


Use # instead of !strconcat( and remove " [$addr];" as this will be added automatically by BasicNVPTXInst. I think this could be simplified by passing ".param" and ".const" for addrspace_name.

AlexMaclean · 2025-07-09T20:05:16Z

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

+defvar frag_pat = (int_nvvm_prefetch_tensormap node:$addr);
+
+def prefetch_tensormap_const : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.const>;
+def prefetch_tensormap_gen : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.generic>;
+def prefetch_tensormap_param : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.param>;


Can you make this part of the multiclass and just pass AS_match.const/generic/param as a parameter?

add prefetch tensormap variant

7e3815b

llvmbot added backend:NVPTX llvm:ir labels Jun 28, 2025

abhilash1910 added 2 commits July 1, 2025 12:04

use generic and const names

775daa3

format

aa8e4d0

durga4github reviewed Jul 1, 2025

View reviewed changes

llvm/docs/NVPTXUsage.rst Outdated Show resolved Hide resolved

durga4github reviewed Jul 1, 2025

View reviewed changes

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td Outdated Show resolved Hide resolved

refresh

6931c80

durga4github reviewed Jul 1, 2025

View reviewed changes

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td Outdated Show resolved Hide resolved

abhilash1910 added 4 commits July 1, 2025 13:39

refactor and refresh

887e139

format

9ddbcfe

format

383d07e

format spaces

0851ae0

durga4github reviewed Jul 1, 2025

View reviewed changes

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td Outdated Show resolved Hide resolved

abhilash1910 added 2 commits July 1, 2025 13:50

refresh

9ea1ca3

refresh

0d26914

durga4github approved these changes Jul 1, 2025

View reviewed changes

durga4github requested a review from AlexMaclean July 1, 2025 11:19

AlexMaclean reviewed Jul 1, 2025

View reviewed changes

llvm/include/llvm/IR/IntrinsicsNVVM.td Outdated Show resolved Hide resolved

abhilash1910 added 3 commits July 4, 2025 15:37

refresh with addrspace

bab560b

update docs

f009291

add isel

2457b1a

abhilash1910 added 4 commits July 7, 2025 15:45

add InferAS

ee0e82c

refresh tests

45f5af0

clang format

ec1e1a0

clang-format

7256bc6

AlexMaclean reviewed Jul 7, 2025

View reviewed changes

abhilash1910 added 2 commits July 7, 2025 23:58

refresh 1

4745db5

refresh 2

1102a58

AlexMaclean reviewed Jul 9, 2025

View reviewed changes

[NVPTX] Add prefetch tensormap variant #146203

Are you sure you want to change the base?

[NVPTX] Add prefetch tensormap variant #146203

Conversation

abhilash1910 commented Jun 28, 2025

Uh oh!

llvmbot commented Jun 28, 2025

Uh oh!

llvmbot commented Jun 28, 2025

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

durga4github left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

github-actions bot commented Jul 7, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

github-actions bot commented Jul 7, 2025 •

edited

Loading